#!/bin/python

#   Copyright 2010 Matthew Hawn

#   Licensed under the Apache License, Version 2.0 (the "License");
#   you may not use this file except in compliance with the License.
#   You may obtain a copy of the License at

#       http://www.apache.org/licenses/LICENSE-2.0

#   Unless required by applicable law or agreed to in writing, software
#   distributed under the License is distributed on an "AS IS" BASIS,
#   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
#   See the License for the specific language governing permissions and
#   limitations under the License.

import pythoncl as cl
import numpy as nd
import array
import struct

platform = cl.get_platforms()[0]
print platform.NAME
devices = platform.get_devices()
print devices[0].NAME
print devices[0].PLATFORM
print devices[0].MAX_WORK_ITEM_SIZES
print devices[0].QUEUE_PROPERTIES
context = devices[0].create_context()
context = cl.create_context(devices=devices)
context = platform.create_context_from_type(device_type=cl.device_type.CPU)
queue = context.create_command_queue(devices[0])
queue.PROFILING_ENABLE = True
buf = context.create_buffer(None, 0, "This is a test")
print queue.enqueue_read_buffer_as_str(buf)
buf = context.create_buffer(None,512)
queue.enqueue_write_buffer(buf,ptr= "This is a test also")
print queue.enqueue_read_buffer_as_str(buf, 17)

m = queue.enqueue_map_buffer(buf,cb=18)
try:
    print buffer(m)
finally:    
    m.unmap().wait() #don't need the wait

m = queue.enqueue_map_buffer(buf, cb=19, map_flags=cl.map_flags.WRITE)
try:
    a = nd.frombuffer(m, dtype=nd.byte)
    print a
    a[3] = ord('g')
finally:
    queue.enqueue_unmap_mem_object(m)
del a
del m

print queue.enqueue_read_buffer_as_str(buf,18)

print "testing build failure:"
program = context.create_program_with_sources("__kerrnel hello() {}")
try:
    program.build()
except cl.cl_error.BUILD_PROGRAM_FAILURE as e:
    print 'Build failed. Looking for info'

    print program.get_build_info(devices[0],cl.program_build_info.LOG)


source = """
__kernel void count(int x, int4 y, constant int* input, global int* output, local int*working) {
    const int id = get_local_id(0);
    int t = get_local_size(0);
    const int mem_start = get_group_id(0) * get_local_size(0) * 2;
    int value;
    int result;
    
    value = working[id] = input[mem_start + id];
    working[id + t] = input[mem_start + id + t];
    barrier(CLK_LOCAL_MEM_FENCE);
    
    while (t>0)
    {
        value += working[id+t];        
        working[id] = value;
        barrier(CLK_LOCAL_MEM_FENCE);
        t >>=1;
    }   
     

     
    if (id == 0) {
        //result = length( (float4) (value * y)) + x;
        result = value;
        //result = t;
        output[get_group_id(0)] = result;
    }
    barrier(CLK_LOCAL_MEM_FENCE);
}
"""

print 'retrying build'
program = context.create_program_with_sources(source)
if not program.build(suppress_error = True):
    print "Unexpected Build Error"
    print program.get_build_info(devices[0],cl.program_build_info.LOG)
print program.get_build_info( devices[0],cl.program_build_info.STATUS)

kernel = program.create_kernel('count')


x = struct.pack('=l', 0)
kernel.set_arg(0,x)

y = struct.pack('=4l', 1,0,0,0)
kernel.set_arg(1,y)

local_size = 32
global_size = 128
item_size = 4
arr = array.array('l', range(global_size))
in_buf = context.create_buffer(cl.mem_flags.READ_ONLY | cl.mem_flags.USE_HOST_PTR, 0, arr)
kernel.set_arg(2,in_buf)

out_buf = context.create_buffer(cl.mem_flags.WRITE_ONLY | cl.mem_flags.ALLOC_HOST_PTR, global_size*item_size/local_size)
kernel.set_arg(3,out_buf)

working = cl.LocalMemory(local_size*item_size)
kernel.set_arg(4, working)

event = queue.enqueue_nd_range_kernel(kernel, 1, None, [global_size/2], [local_size/2])
event.wait()

#arr = array.array('l', [0] * (global_size/local_size))
#print arr.tolist()
#queue.enqueue_read_buffer(out_buf, arr)
s = queue.enqueue_read_buffer_as_str(out_buf)

#print arr.tolist()
print struct.unpack('=4l', s)






